#include "LinearInterpolation.h"
#include <thrust\fill.h>
#include <thrust\device_ptr.h>




__global__ void interpolate_kernel(float* device_in, float* device_out, int width, int height){
	int x = blockIdx.x * blockDim.x + threadIdx.x;
	int y = blockIdx.y * blockDim.y + threadIdx.y;

	int out_x = x * 2;
	int out_y = y * 2;
	
	device_out[out_x + out_y * 640*2] = device_in[x + y * 640];
	
	/*device_out[(out_x + 1) + out_y * (width * 2)] = 0.0;
	device_out[out_x + (out_y + 1) * (width * 2)] = 0.0;*/

	if(device_in[x + y * width] > 50){	
		if((x + 1) + y * width < width * height && device_in[(x + 1) + y * width] > 50)
			device_out[(out_x + 1) + out_y * (width * 2)] = 
				(device_in[x + y * width] + device_in[(x + 1) + y * width]) / 2.0f;

		if(x + (y + 1) * width < width * height && device_in[x + (y + 1) * width] > 50)
			device_out[out_x + (out_y + 1) * (width * 2)] =
				(device_in[x + y * width] + device_in[x + (y + 1) * width]) / 2.0f;
		
		if((x + 1) + (y + 1) * width < width * height && device_in[(x + 1) + (y + 1) * width] > 50)
			device_out[(out_x + 1) + (out_y + 1) * (width * 2)] =
				(device_in[x + y * width] + device_in[(x + 1) + (y + 1) * width]) / 2.0f;
	} 
}





float* LinearInterpolation::interpolate(float* in){
	thrust::fill(
		thrust::device_ptr<float>(out_device),
		thrust::device_ptr<float>(out_device) + 1280*960,
		0.0
		);


	//cudaMemcpy(in_device, in, sizeof(float) * w * h, cudaMemcpyDeviceToDevice);

	interpolate_kernel<<<dim3(20, 20), dim3(32, 24)>>>
		(in, out_device, w, h);

	//cudaMemcpy(out_host, out_device, sizeof(float) * w * 2 * h * 2, cudaMemcpyDeviceToHost);
	//display();
	return out_device;
}